//
//  convolute_kernel.cl
//  FilterModule_Standalone
//
//  Created by kadlubek47 on 28.04.2013.
//  Copyright (c) 2013 kadlubek47. All rights reserved.
//



//----------------------------------------------------
//one channel
__kernel
void convoluteChannel( __global float* x_global,
                       __local float* x,
                       __constant float* h,
                       __constant unsigned int* M,
                       __global float* y
                       )
{
    int index = get_global_id(0);
    
    y[index] = 0;
    int N = 2 * (*M) + 1;
    for(int i = 0; i < N; ++i) {
       y[index] += x_global[index + i] * h[i];
    }
}
//----------------------------------------------------



//----------------------------------------------------
//multichannel without local memory optimization
__kernel
void convoluteChannels(__global float* inputChannelsWindowLinear,
                       __global float* outputChannelsLinear,
                       __global float* ir,
                       __global unsigned int* infoTab //infoTab[0] == channelsLength, infoTab[1] == irLength
                       )
{
    unsigned int index = get_global_id(0);
    unsigned int channNum  = index / infoTab[0];
    unsigned int sampleNum = index % infoTab[0];
    
    outputChannelsLinear[index] = 0; //[channNum][sampleNum]
    unsigned int baseInputIndex = channNum * infoTab[0] + sampleNum;
    for(unsigned int i = 0; i < infoTab[1]; ++i)
        outputChannelsLinear[index] += inputChannelsWindowLinear[baseInputIndex + i] * ir[i];
}
//----------------------------------------------------


